//----------------------------------------------------------------------------
// M A R I T I M E  R E S E A R C H  I N S T I T U T E  N E T H E R L A N D S
//----------------------------------------------------------------------------
// Copyright (C) 2012 - MARIN - All rights reserved - http://www.marin.nl
//----------------------------------------------------------------------------
// Program    : mermaid
// Module     : vbm
// File       : subs.h
// Author     : M.A. de Jong
//----------------------------------------------------------------------------

#ifndef SUBS_H
#define	SUBS_H

#include "kernels/subs_kernels.h"

namespace rrbsolver
{

template <class T>
inline void subs_forward(T* x, const Grid orig)
{
    dimGrid.x = orig.cx / NUM_THREADS_X;
    dimGrid.y = orig.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, x, &desc, 
      orig.nx, orig.ny, orig.nx * sizeof(T)), 
        "rrbsolver::texture bind XX @ kernel_subs1");
    checkCudaError(cudaBindTexture2D(0, &texCC, orig.cc, &desc, 
      orig.nx, orig.ny, orig.nx * sizeof(T)), 
        "rrbsolver::texture bind CC @ kernel_subs1");
#endif

    kernel_subs1<T> <<< dimGrid, dimBlock >>> (x, orig);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_subs1");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture XX @ kernel_subs1");
    checkCudaError(cudaUnbindTexture(texCC), 
        "rrbsolver::unbind texture CC @ kernel_subs1");
#endif
}

template <class T>
inline void subs_backward(T* x, const T* r, const Grid orig)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, x, &desc, 
      orig.nx, orig.ny, orig.nx * sizeof(T)), 
        "rrbsolver::texture bind XX @ kernel_subs2");
    checkCudaError(cudaBindTexture2D(0, &texCC, orig.cc, &desc, 
      orig.nx, orig.ny, orig.nx * sizeof(T)), 
        "rrbsolver::texture bind CC @ kernel_subs2");
#endif

    dimGrid.x = orig.cx / NUM_THREADS_X;
    dimGrid.y = orig.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_subs2<T> <<< dimGrid, dimBlock >>> (x, r, orig);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_subs2");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture XX @ kernel_subs2");
    checkCudaError(cudaUnbindTexture(texCC), 
        "rrbsolver::unbind texture CC @ kernel_subs2");
#endif
}

} // namespace rrbsolver

#endif // SUBS_H

